/*
 * This file is part of AdaptiveCpp, an implementation of SYCL and C++ standard
 * parallelism for CPUs and GPUs.
 *
 * Copyright The AdaptiveCpp Contributors
 *
 * AdaptiveCpp is released under the BSD 2-Clause "Simplified" License.
 * See file LICENSE in the project root for full license details.
 */
// SPDX-License-Identifier: BSD-2-Clause
#ifndef HIPSYCL_ALGORITHMS_ALGORITHM_HPP
#define HIPSYCL_ALGORITHMS_ALGORITHM_HPP

#include <functional>
#include <iterator>
#include <limits>
#include <type_traits>
#include <cstring>
#include "hipSYCL/sycl/libkernel/accessor.hpp"
#include "hipSYCL/sycl/libkernel/atomic_builtins.hpp"
#include "hipSYCL/sycl/libkernel/memory.hpp"
#include "hipSYCL/sycl/libkernel/functional.hpp"
#include "hipSYCL/sycl/detail/namespace_compat.hpp"
#include "hipSYCL/sycl/event.hpp"
#include "hipSYCL/sycl/queue.hpp"
#include "merge/merge.hpp"
#include "scan/scan.hpp"
#include "util/traits.hpp"
#include "hipSYCL/algorithms/numeric.hpp"
#include "hipSYCL/algorithms/util/allocation_cache.hpp"
#include "hipSYCL/algorithms/util/memory_streaming.hpp"
#include "hipSYCL/algorithms/sort/bitonic_sort.hpp"
#include "hipSYCL/algorithms/merge/merge.hpp"
#include "hipSYCL/algorithms/scan/scan.hpp"


namespace hipsycl::algorithms {

namespace detail {

template<class T>
bool all_bytes_equal(const T& val, unsigned char& byte_value) {
  std::array<unsigned char, sizeof(T)> buff;
  std::memcpy(buff.data(), &val, sizeof(T));

  for(int i = 0; i < sizeof(T); ++i) {
    if(buff[i] != buff[0])
      return false;
  }
  byte_value = buff[0];
  return true;
}

inline bool should_use_memcpy(const sycl::device& dev) {
#ifdef __ACPP_STDPAR_ENABLE_AUTO_MULTIQUEUE__
  // MQS assumes that all operations are kernels
  return false;
#endif

  // OpenMP backend implements queue::memcpy() using std::memcpy
  // which can break perf on NUMA systems
  if(dev.get_backend() == sycl::backend::omp)
    return false;
  // Some OpenCL implementations (e.g. Intel GPU) seem to be very
  // inefficient for memcpy calls between two pointers if data
  // is source and dest are on the same device (does it always go
  // through the host?)
  if(dev.get_backend() == sycl::backend::ocl)
    return false;
  if(dev.get_backend() == sycl::backend::hip)
    // It was reported that hipMemcpy does not handly copies involving
    // shared allocations efficiently
    return false;

  return true;
}

inline bool should_use_memset(const sycl::device& dev) {
#ifdef __ACPP_STDPAR_ENABLE_AUTO_MULTIQUEUE__
  // MQS assumes that all operations are kernels
  return false;
#endif

  if(dev.get_backend() == sycl::backend::omp)
    return false;
  if(dev.get_backend() == sycl::backend::ocl)
    return false;
  if(dev.get_backend() == sycl::backend::hip)
    return false;

  return true;
}

}

template <class ForwardIt, class UnaryFunction2>
sycl::event for_each(sycl::queue &q, ForwardIt first, ForwardIt last,
                     UnaryFunction2 f,
                     const std::vector<sycl::event> &deps = {}) {
  if(first == last)
    return sycl::event{};
  return q.parallel_for(sycl::range{std::distance(first, last)}, deps,
                        [=](sycl::id<1> id) {
                          auto it = first;
                          std::advance(it, id[0]);
                          f(*it);
                        });
}

template <class ForwardIt, class Size, class UnaryFunction2>
sycl::event for_each_n(sycl::queue &q, ForwardIt first, Size n,
                       UnaryFunction2 f,
                       const std::vector<sycl::event> &deps = {}) {
  if(n <= 0)
    // sycl::event{} represents a no-op that is always finished.
    // This means it does not respect prior tasks in the task graph!
    // TODO Is this okay? Can we defer this responsibility to the user?
    return sycl::event{};
  return q.parallel_for(sycl::range{static_cast<size_t>(n)}, deps,
                        [=](sycl::id<1> id) {
                          auto it = first;
                          std::advance(it, id[0]);
                          f(*it);
                        });
}

template <class ForwardIt1, class ForwardIt2, class UnaryOperation>
sycl::event transform(sycl::queue &q, ForwardIt1 first1, ForwardIt1 last1,
                      ForwardIt2 d_first, UnaryOperation unary_op,
                      const std::vector<sycl::event> &deps = {}) {
  if(first1 == last1)
    return sycl::event{};
  return q.parallel_for(sycl::range{std::distance(first1, last1)}, deps,
                        [=](sycl::id<1> id) {
                          auto input = first1;
                          auto output = d_first;
                          std::advance(input, id[0]);
                          std::advance(output, id[0]);
                          *output = unary_op(*input);
                        });
}

template <class ForwardIt1, class ForwardIt2, class ForwardIt3,
          class BinaryOperation>
sycl::event transform(sycl::queue &q, ForwardIt1 first1, ForwardIt1 last1,
                      ForwardIt2 first2, ForwardIt3 d_first,
                      BinaryOperation binary_op,
                      const std::vector<sycl::event> &deps = {}) {
  if(first1 == last1)
    return sycl::event{};
  return q.parallel_for(sycl::range{std::distance(first1, last1)}, deps,
                        [=](sycl::id<1> id) {
                          auto input1 = first1;
                          auto input2 = first2;
                          auto output = d_first;
                          std::advance(input1, id[0]);
                          std::advance(input2, id[0]);
                          std::advance(output, id[0]);
                          *output = binary_op(*input1, *input2);
                        });
}

template <class ForwardIt1, class ForwardIt2>
sycl::event copy(sycl::queue &q, ForwardIt1 first, ForwardIt1 last,
                 ForwardIt2 d_first, const std::vector<sycl::event> &deps = {}) {
  
  auto size = std::distance(first, last);
  if(size == 0)
    return sycl::event{};
  
  using value_type1 = typename std::iterator_traits<ForwardIt1>::value_type;
  using value_type2 = typename std::iterator_traits<ForwardIt2>::value_type;

  if (std::is_trivially_copyable_v<value_type1> &&
      std::is_same_v<value_type1, value_type2> &&
      util::is_contiguous<ForwardIt1>() && util::is_contiguous<ForwardIt2>() &&
      detail::should_use_memcpy(q.get_device())) {
    return q.memcpy(&(*d_first), &(*first), size * sizeof(value_type1), deps);
  } else {
    return q.parallel_for(sycl::range{size}, deps,
                          [=](sycl::id<1> id) {
                            auto input = first;
                            auto output = d_first;
                            std::advance(input, id[0]);
                            std::advance(output, id[0]);
                            *output = *input;
                          });
  }
}

template <class ForwardIt1, class ForwardIt2, class UnaryPredicate>
sycl::event copy_if(sycl::queue &q, util::allocation_group &scratch_allocations,
                    ForwardIt1 first, ForwardIt1 last, ForwardIt2 d_first,
                    UnaryPredicate pred,
                    std::size_t *num_elements_copied = nullptr,
                    const std::vector<sycl::event> &deps = {}) {
  if(first == last) {
    if(num_elements_copied)
      *num_elements_copied = 0;
    return sycl::event{};
  }

  // TODO: We could optimize by switching between 32/64 bit types
  // depending on problem size
  using ScanT = std::size_t;

  auto generator = [=](auto idx, auto effective_group_id,
                       auto effective_global_id, auto problem_size) {
    if(effective_global_id >= problem_size)
      return ScanT{0};

    ForwardIt1 it = first;
    std::advance(it, effective_global_id);
    if(pred(*it))
      return ScanT{1};

    return ScanT{0};
  };

  auto result_processor = [=](auto idx, auto effective_group_id,
                       auto effective_global_id, auto problem_size,
                       auto value) {
    if (effective_global_id < problem_size) {
      ForwardIt2 output = d_first;
      ForwardIt1 input = first;
      std::advance(input, effective_global_id);
      std::advance(output, value);

      bool needs_copy = false;

      if(effective_global_id < problem_size) {
        auto input_value = *input;
        needs_copy = pred(input_value);
        if(needs_copy)
          *output = *input;
      }

      if (effective_global_id == problem_size - 1 && num_elements_copied) {
        ScanT inclusive_scan_result = value;
        // We did an exclusive scan, so if the last element also was copied,
        // we need to add that.
        if(needs_copy)
          ++inclusive_scan_result;
        
        *num_elements_copied = static_cast<std::size_t>(inclusive_scan_result);
      }
    }
  };

  std::size_t problem_size = std::distance(first, last);

  constexpr bool is_inclusive_scan = false;
  return scanning::generate_scan_process<is_inclusive_scan, ScanT>(
      q, scratch_allocations, problem_size, sycl::plus<>{},
      ScanT{0}, generator, result_processor, deps);
}

template <class ForwardIt1, class Size, class ForwardIt2>
sycl::event copy_n(sycl::queue &q, ForwardIt1 first, Size count,
                   ForwardIt2 result,
                   const std::vector<sycl::event> &deps = {}) {
  if(count <= 0)
    return sycl::event{};

  auto last = first;
  std::advance(last, count);
  return copy(q, first, last, result, deps);
}

template <class ForwardIt1, class ForwardIt2>
sycl::event move(sycl::queue &q, ForwardIt1 first, ForwardIt1 last,
                 ForwardIt2 d_first, const std::vector<sycl::event> &deps = {}) {

  auto size = std::distance(first, last);
  if (size == 0)
    return sycl::event{};

  using value_type1 = typename std::iterator_traits<ForwardIt1>::value_type;
  using value_type2 = typename std::iterator_traits<ForwardIt2>::value_type;

  if (std::is_trivially_copyable_v<value_type1> &&
      std::is_same_v<value_type1, value_type2> &&
      util::is_contiguous<ForwardIt1>() && util::is_contiguous<ForwardIt2>() &&
      detail::should_use_memcpy(q.get_device())) {
    return q.memcpy(&(*d_first), &(*first), size * sizeof(value_type1), deps);
  } else {
    return q.parallel_for(sycl::range{size}, deps,
                          [=](sycl::id<1> id) {
                            auto input = first;
                            auto output = d_first;
                            std::advance(input, id[0]);
                            std::advance(output, id[0]);
                            *output = std::move(*input);
                          });
  }
}

template <class ForwardIt, class T>
sycl::event fill(sycl::queue &q, ForwardIt first, ForwardIt last,
                 const T &value, const std::vector<sycl::event> &deps = {}) {
  auto size = std::distance(first, last);
  if(size == 0)
    return sycl::event{};

  using value_type = typename std::iterator_traits<ForwardIt>::value_type;

  auto invoke_kernel = [&]() -> sycl::event{
    return q.parallel_for(sycl::range{size}, deps,
                        [=](sycl::id<1> id) {
                          auto it = first;
                          std::advance(it, id[0]);
                          *it = value;
                        });
  };

  if constexpr (std::is_trivial_v<value_type> &&
                std::is_same_v<value_type, T> &&
                util::is_contiguous<ForwardIt>()) {
    unsigned char equal_byte;
    if (detail::all_bytes_equal(value, equal_byte) &&
        detail::should_use_memset(q.get_device())) {
      return q.memset(&(*first), static_cast<int>(equal_byte),
                      size * sizeof(T), deps);
    } else {
      return invoke_kernel();
    }
  } else {
    return invoke_kernel();
  }
}

template<class ForwardIt, class Size, class T >
sycl::event fill_n(sycl::queue& q,
                  ForwardIt first, Size count, const T& value,
                  const std::vector<sycl::event> &deps = {}) {
  if(count <= Size{0})
    return sycl::event{};
  
  auto last = first;
  std::advance(last, count);
  return fill(q, first, last, value, deps);
}

template <class ForwardIt, class Generator>
sycl::event generate(sycl::queue &q, ForwardIt first, ForwardIt last,
                     Generator g, const std::vector<sycl::event> &deps = {}) {
  if(first == last)
    return sycl::event{};
  return q.parallel_for(sycl::range{std::distance(first, last)}, deps,
                        [=](sycl::id<1> id) {
                          auto it = first;
                          std::advance(it, id[0]);
                          *it = g();
                        });
}

template <class ForwardIt, class Size, class Generator>
sycl::event generate_n(sycl::queue &q, ForwardIt first, Size count, Generator g,
                       const std::vector<sycl::event> &deps = {}) {
  if(count <= 0)
    return sycl::event{};
  return q.parallel_for(sycl::range{static_cast<size_t>(count)}, deps,
                        [=](sycl::id<1> id) {
                          auto it = first;
                          std::advance(it, id[0]);
                          *it = g();
                        });
}

template <class ForwardIt1, class ForwardIt2, class T>
sycl::event remove_copy(sycl::queue &q, util::allocation_group &scratch_allocations,
                    ForwardIt1 first, ForwardIt1 last, ForwardIt2 d_first,
                    const T &value, std::size_t *num_elements_copied = nullptr,
                    const std::vector<sycl::event> &deps = {}) {
    auto pred = [value](auto x){return !(x == value);};
    return copy_if(q, scratch_allocations, first, last, d_first, pred,
                   num_elements_copied, deps);
}

template <class ForwardIt, class T>
sycl::event remove(sycl::queue &q, util::allocation_group &scratch_allocations,
                   ForwardIt first, ForwardIt last, const T &value,
                   std::size_t *num_elements_copied = nullptr,
                   const std::vector<sycl::event> &deps = {}) {
  if(first == last) {
    if(num_elements_copied)
      *num_elements_copied = 0;
    return sycl::event{};
  }

  T* device_buffer = scratch_allocations.obtain<T>(std::distance(first, last));

  auto pred = [value](auto x){ return !(x == value); };
  auto evt = copy_if(q, scratch_allocations, first, last, device_buffer, pred,
          num_elements_copied, deps);

  evt.wait();
  return copy_n(q, &(*device_buffer), *num_elements_copied, first,
                deps);
}

template <class ForwardIt, class UnaryPredicate>
sycl::event remove_if(sycl::queue &q, util::allocation_group &scratch_allocations,
                   ForwardIt first, ForwardIt last, UnaryPredicate pred,
                   std::size_t *num_elements_copied = nullptr,
                   const std::vector<sycl::event> &deps = {}) {
  if(first == last) {
    if(num_elements_copied)
      *num_elements_copied = 0;
    return sycl::event{};
  }

  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;
  ValueT* device_buffer = scratch_allocations.obtain<ValueT>(std::distance(first, last));

  auto op = [pred](auto x){ return !pred(x); };
  auto evt = copy_if(q, scratch_allocations, first, last, device_buffer, op,
          num_elements_copied, deps);

  evt.wait();
  return copy_n(q, &(*device_buffer), *num_elements_copied, first,
                deps);
}

template <class ForwardIt1, class ForwardIt2, class UnaryPredicate>
sycl::event remove_copy_if(sycl::queue &q, util::allocation_group &scratch_allocations,
                    ForwardIt1 first, ForwardIt1 last, ForwardIt2 d_first,
                    UnaryPredicate p, std::size_t *num_elements_copied = nullptr,
                    const std::vector<sycl::event> &deps = {}) {
    auto pred = [p](auto x){return !p(x);};
    return copy_if(q, scratch_allocations, first, last, d_first, pred,
                   num_elements_copied, deps);
}

template <class ForwardIt, class T>
sycl::event replace(sycl::queue &q, ForwardIt first, ForwardIt last,
                    const T &old_value, const T &new_value,
                    const std::vector<sycl::event> &deps = {}) {
  if(first == last)
    return sycl::event{};
  return for_each(q, first, last,[=](auto& x){
    if(x == old_value)
      x = new_value;
  }, deps);
}

template <class ForwardIt, class UnaryPredicate, class T>
sycl::event replace_if(sycl::queue &q, ForwardIt first, ForwardIt last,
                       UnaryPredicate p, const T &new_value,
                       const std::vector<sycl::event> &deps = {}) {
  if(first == last)
    return sycl::event{};
  return for_each(q, first, last, [=](auto& x){
    if(p(x))
      x = new_value;
  }, deps);
}

template <class ForwardIt1, class ForwardIt2, class UnaryPredicate, class T>
sycl::event replace_copy_if(sycl::queue &q, ForwardIt1 first, ForwardIt1 last,
                            ForwardIt2 d_first, UnaryPredicate p,
                            const T &new_value,
                            const std::vector<sycl::event> &deps = {}) {
  if (first == last)
    return sycl::event{};
  return q.parallel_for(sycl::range{std::distance(first, last)}, deps,
                        [=](sycl::id<1> id) {
                          auto input = first;
                          auto output = d_first;
                          std::advance(input, id[0]);
                          std::advance(output, id[0]);
                          if (p(*input)) {
                            *output = new_value;
                          } else {
                            *output = *input;
                          }
                        });
}

template <class ForwardIt1, class ForwardIt2, class T>
sycl::event replace_copy(sycl::queue &q, ForwardIt1 first, ForwardIt1 last,
                         ForwardIt2 d_first, const T &old_value,
                         const T &new_value,
                         const std::vector<sycl::event> &deps = {}) {
  if (first == last)
    return sycl::event{};
  return replace_copy_if(
      q, first, last, d_first, [=](const auto &x) { return x == old_value; },
      new_value, deps);
}

template <class BidirIt>
sycl::event reverse(sycl::queue &q, BidirIt first, BidirIt last,
                     const std::vector<sycl::event> &deps = {}) {
  auto size = std::distance(first, last);
  if (first == last || size == 1)
    return sycl::event{};

  return q.parallel_for(sycl::range{size/2}, deps,
                        [=](sycl::id<1> id) {
                          auto offset = size - id[0] - 1;
                          auto input = std::next(first, id[0]);
                          auto output = std::next(first, offset);
                          std::iter_swap(input, output);
                        });
}

template <class BidirIt, class ForwardIt>
sycl::event reverse_copy(sycl::queue &q, BidirIt first,
                         BidirIt last, ForwardIt d_first,
                         const std::vector<sycl::event> &deps = {}) {
  if (first == last)
    return sycl::event{};

  auto size = std::distance(first, last);

  return q.parallel_for(sycl::range{size}, deps,
                        [=](sycl::id<1> id) {
                          auto offset = size - id[0] - 1;
                          auto input = std::next(first, offset);
                          auto output = std::next(d_first, id[0]);
                          *output = *input;
                        });
}

template <class ForwardIt, class T>
sycl::event find(sycl::queue &q, util::allocation_group &scratch_allocations,
                 ForwardIt first, ForwardIt last, const T &value,
                 typename std::iterator_traits<ForwardIt>::difference_type *out,
                 const std::vector<sycl::event> &deps = {}) {
  if(first == last)
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [first, value, problem_size] (ForwardIt input) {
    return (*input == value ? std::distance(first, input) : problem_size);
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    reducer.combine(transform(input));
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt, class UnaryPredicate>
sycl::event find_if(sycl::queue &q, util::allocation_group &scratch_allocations,
                 ForwardIt first, ForwardIt last, UnaryPredicate p,
                 typename std::iterator_traits<ForwardIt>::difference_type *out,
                 const std::vector<sycl::event> &deps = {}) {
  if(first == last)
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [first, p, problem_size] (ForwardIt input) {
    return (p(*input) ? std::distance(first, input) : problem_size);
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    // Transform returns a DiffT object.
    reducer.combine(transform(input));
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}


template <class ForwardIt, class UnaryPredicate>
sycl::event find_if_not(sycl::queue &q, util::allocation_group &scratch_allocations,
                 ForwardIt first, ForwardIt last, UnaryPredicate p,
                 typename std::iterator_traits<ForwardIt>::difference_type *out,
                 const std::vector<sycl::event> &deps = {}) {
  if(first == last)
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [first, p, problem_size] (ForwardIt input) {
    return (!p(*input) ? std::distance(first, input) : problem_size);
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    // Transform returns a DiffT object.
    reducer.combine(transform(input));
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt1, class ForwardIt2>
sycl::event find_end(sycl::queue &q, util::allocation_group &scratch_allocations,
                          ForwardIt1 first, ForwardIt1 last, ForwardIt2 s_first,
                          ForwardIt2 s_last,
                          typename std::iterator_traits<ForwardIt1>::difference_type *out,
                          const std::vector<sycl::event> &deps = {}) {
  if (first == last || s_first == s_last)
    return sycl::event{};

  if (std::distance(first, last) < std::distance(s_first, s_last))
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt1>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [=](ForwardIt1 input) {
    auto it = input;
    for (ForwardIt1 s_it = s_first; s_it != s_last; ++s_it) {
      if(!(*it++ == *s_it))
        return std::numeric_limits<DiffT>::min();
    }
    return std::distance(first, input);
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    reducer.combine(transform(input));
  };

  auto reduce = sycl::maximum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::min(),
                                       problem_size, kernel, reduce, deps);
}


template <class ForwardIt1, class ForwardIt2, class BinaryPredicate>
sycl::event find_end(sycl::queue &q, util::allocation_group &scratch_allocations,
                          ForwardIt1 first, ForwardIt1 last, ForwardIt2 s_first,
                          ForwardIt2 s_last, BinaryPredicate p,
                          typename std::iterator_traits<ForwardIt1>::difference_type *out,
                          const std::vector<sycl::event> &deps = {}) {
  if (first == last || s_first == s_last)
    return sycl::event{};

  if (std::distance(first, last) < std::distance(s_first, s_last))
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt1>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [=](ForwardIt1 input) {
    auto it = input;
    for (ForwardIt1 s_it = s_first; s_it != s_last; ++s_it) {
      if( !p(*it++, *s_it) )
        return std::numeric_limits<DiffT>::min();
    }
    return std::distance(first, input);
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    reducer.combine(transform(input));
  };

  auto reduce = sycl::maximum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::min(),
                                       problem_size, kernel, reduce, deps);
}


template <class ForwardIt1, class ForwardIt2>
sycl::event find_first_of(sycl::queue &q, util::allocation_group &scratch_allocations,
                          ForwardIt1 first, ForwardIt1 last, ForwardIt2 s_first,
                          ForwardIt2 s_last,
                          typename std::iterator_traits<ForwardIt1>::difference_type *out,
                          const std::vector<sycl::event> &deps = {}) {
  if (first == last || s_first == s_last)
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt1>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [=] (ForwardIt1 input) {
    auto s_it = s_first;
    for (; s_it != s_last; ++s_it) {
      if(*input == *s_it)
        return std::distance(first, input);
    }
    return problem_size;
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    reducer.combine(transform(input));
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt1, class ForwardIt2, class BinaryPredicate>
sycl::event find_first_of(sycl::queue &q, util::allocation_group &scratch_allocations,
                          ForwardIt1 first, ForwardIt1 last, ForwardIt2 s_first,
                          ForwardIt2 s_last, BinaryPredicate p,
                          typename std::iterator_traits<ForwardIt1>::difference_type *out,
                          const std::vector<sycl::event> &deps = {}) {
  if (first == last || s_first == s_last)
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt1>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [=] (ForwardIt1 input) {
    auto s_it = s_first;
    for (; s_it != s_last; ++s_it) {
      if(p(*input, *s_it))
        return std::distance(first, input);
    }
    return problem_size;
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    reducer.combine(transform(input));
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}

// Need transform_reduce functionality for find etc, so forward
// declare here.
/*template <class ForwardIt, class T, class BinaryReductionOp,
          class UnaryTransformOp>
sycl::event
transform_reduce(sycl::queue &q, util::allocation_group &scratch_allocations,
                 ForwardIt first, ForwardIt last, T* out, T init,
                 BinaryReductionOp reduce, UnaryTransformOp transform,
                 const std::vector<sycl::event> &deps);

// Need transform_reduce functionality for find etc, so forward
// declare here.
template <class ForwardIt, class T, class BinaryReductionOp,
          class UnaryTransformOp>
sycl::event
transform_reduce(sycl::queue &q, util::allocation_group &scratch_allocations,
                 ForwardIt first, ForwardIt last, T* out, T init,
                 BinaryReductionOp reduce, UnaryTransformOp transform);
*/

namespace detail {
using early_exit_flag_t = int;

// predicate must be a callable of type bool(sycl::id<1>).
// If it returns true, the for_each will abort and output_has_exited_early
// will be set to true.
template <class Predicate>
sycl::event early_exit_for_each(sycl::queue &q, std::size_t problem_size,
                                early_exit_flag_t *output_has_exited_early,
                                Predicate should_exit,
                                const std::vector<sycl::event> &deps = {}) {
  
  std::size_t group_size = 128;

  util::abortable_data_streamer streamer{q.get_device(), problem_size, group_size};

  std::size_t dispatched_global_size = streamer.get_required_global_size();

  auto kernel = [=](sycl::nd_item<1> idx) {
    util::abortable_data_streamer::run(
        problem_size, idx, output_has_exited_early,
        [&](sycl::id<1> idx) { return should_exit(idx); });
  };

  auto evt = q.single_task(deps, [=](){*output_has_exited_early = false;});
  return q.parallel_for(sycl::nd_range<1>{dispatched_global_size, group_size}, evt,
                        kernel);
}

}

template <class ForwardIt, class UnaryPredicate>
sycl::event all_of(sycl::queue &q,
                   ForwardIt first, ForwardIt last, detail::early_exit_flag_t* out,
                   UnaryPredicate p, const std::vector<sycl::event>& deps = {}) {
  std::size_t problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};
  auto evt = detail::early_exit_for_each(q, problem_size, out,
                                     [=](sycl::id<1> idx) -> bool {
                                       auto it = first;
                                       std::advance(it, idx[0]);
                                       return !p(*it);
                                     }, deps);
  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template <class ForwardIt, class UnaryPredicate>
sycl::event any_of(sycl::queue &q,
                   ForwardIt first, ForwardIt last, detail::early_exit_flag_t* out,
                   UnaryPredicate p, const std::vector<sycl::event>& deps = {}) {
  std::size_t problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};
  return detail::early_exit_for_each(q, problem_size, out,
                                     [=](sycl::id<1> idx) -> bool {
                                       auto it = first;
                                       std::advance(it, idx[0]);
                                       return p(*it);
                                     }, deps);
}

template <class ForwardIt, class UnaryPredicate>
sycl::event none_of(sycl::queue &q,
                   ForwardIt first, ForwardIt last, detail::early_exit_flag_t* out,
                   UnaryPredicate p, const std::vector<sycl::event>& deps = {}) {
  std::size_t problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};
  
  auto evt = any_of(q, first, last, out, p, deps);
  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template<class ForwardIt, class T>
sycl::event count(sycl::queue &q, util::allocation_group &scratch_allocations,
                  ForwardIt first, ForwardIt last,
                  typename std::iterator_traits<ForwardIt>::difference_type *out,
                  const T& value, const std::vector<sycl::event> &deps = {}) {

  using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;

  return transform_reduce(q, scratch_allocations, first, last, out,
                          DiffT{}, std::plus<>{},
                          [value](ValueT x) {return (x == value ? 1 : 0);},
                          deps);
}

template<class ForwardIt, class UnaryPredicate>
sycl::event count_if(sycl::queue &q, util::allocation_group &scratch_allocations,
                  ForwardIt first, ForwardIt last,
                  typename std::iterator_traits<ForwardIt>::difference_type *out,
                  UnaryPredicate p, const std::vector<sycl::event> &deps = {}) {

  using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;

  return transform_reduce(q, scratch_allocations, first, last, out,
                          DiffT{}, std::plus<>{},
                          [p](ValueT x) {return p(x) ? 1 : 0;},
                          deps);
}

template <class ForwardIt1, class ForwardIt2, class BinaryPredicate>
sycl::event mismatch(sycl::queue &q, util::allocation_group &scratch_allocations,
                    ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                    ForwardIt2 last2, BinaryPredicate p,
                    typename std::iterator_traits<ForwardIt1>::difference_type* out,
                    const std::vector<sycl::event>& deps = {}) {
  if (first1 == last1 || first2 == last2)
    return sycl::event{};

  using  DiffT = typename std::iterator_traits<ForwardIt1>::difference_type;
  DiffT problem_size = std::min(std::distance(first1, last1),
                                std::distance(first2, last2));

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input1 = std::next(first1, idx[0]);
    auto input2 = std::next(first2, idx[0]);
    if ( p(*input1, *input2) )
      reducer.combine(problem_size);
    else
      reducer.combine(idx[0]);
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out,
                                       std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt1, class ForwardIt2, class BinaryPredicate>
sycl::event mismatch(sycl::queue &q, util::allocation_group &scratch_allocations,
                    ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                    BinaryPredicate p,
                    typename std::iterator_traits<ForwardIt1>::difference_type* out,
                    const std::vector<sycl::event>& deps = {}) {
  if (first1 == last1)
    return sycl::event{};

  return mismatch(q, scratch_allocations, first1, last1, first2,
                  std::next(first2, std::distance(first1, last1)), p, out, deps);
}

template <class ForwardIt1, class ForwardIt2>
sycl::event mismatch(sycl::queue &q, util::allocation_group &scratch_allocations,
                    ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                    ForwardIt2 last2,
                    typename std::iterator_traits<ForwardIt1>::difference_type* out,
                    const std::vector<sycl::event>& deps = {}) {
  if (first1 == last1 || first2 == last2)
    return sycl::event{};

  return mismatch(q, scratch_allocations, first1, last1, first2,
                  last2, std::equal_to<>(), out, deps);
}

template <class ForwardIt1, class ForwardIt2>
sycl::event mismatch(sycl::queue &q, util::allocation_group &scratch_allocations,
                    ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                    typename std::iterator_traits<ForwardIt1>::difference_type* out,
                    const std::vector<sycl::event>& deps = {}) {
  if (first1 == last1)
    return sycl::event{};

  return mismatch(q, scratch_allocations, first1, last1, first2,
                  std::next(first2, std::distance(first1, last1)), out, deps);
}

template <class ForwardIt>
sycl::event
min_element(sycl::queue &q, util::allocation_group &scratch_allocations,
            ForwardIt first, ForwardIt last,
            std::pair<ForwardIt, typename
            std::iterator_traits<ForwardIt>::value_type> *out,
            const std::vector<sycl::event> &deps= {}) {
  auto problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};

  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;
  using MinPair = std::pair<ForwardIt, ValueT>;

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    MinPair p = std::make_pair(input, *input);
    reducer.combine(p);
  };

  auto reduce = [first] (MinPair a, MinPair b) {
    // Preserve strict total order over two equivalent
    // pointers, i.e. return the element that appears
    // in the sequence nearest to first.
    if (!(a.second < b.second) && !(b.second < a.second)) {
      if (std::distance(first, a.first) < std::distance(first, b.first))
        return a;
      else
        return b;
    }
#if __cplusplus < 202002L
    else if (a.second < b.second)
#else
    else if (std::less{}(a.second, b.second))
#endif
      return a;
    else
      return b;
  };

  MinPair init = std::make_pair(first, *first);

  return detail::transform_reduce_impl(q, scratch_allocations, out, init,
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt, class Compare>
sycl::event
min_element(sycl::queue &q, util::allocation_group &scratch_allocations,
            ForwardIt first, ForwardIt last, Compare comp,
            std::pair<ForwardIt, typename
            std::iterator_traits<ForwardIt>::value_type> *out,
            const std::vector<sycl::event> &deps= {}) {
  auto problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};

  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;
  using MinPair = std::pair<ForwardIt, ValueT>;

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    MinPair p = std::make_pair(input, *input);
    reducer.combine(p);
  };

  auto reduce = [comp, first] (MinPair a, MinPair b) {
    // Comp used for associative containers must always
    // return false for equal values. (Effective STL, Item 21)
    // In cases where it does not (for eg. std::less_equal),
    // implementation aligns behaviour to libstdc++, returning
    // the element furthest from first.
    if (std::distance(first, a.first) < std::distance(first, b.first))
      if (comp(b.second, a.second) == false)
        return a;
      else
        return b;
    else
      if (comp(a.second, b.second) == false)
        return b;
      else
        return a;
  };

  MinPair init = std::make_pair(first, *first);

  return detail::transform_reduce_impl(q, scratch_allocations, out, init,
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt>
sycl::event
max_element(sycl::queue &q, util::allocation_group &scratch_allocations,
            ForwardIt first, ForwardIt last,
            std::pair<ForwardIt, typename
            std::iterator_traits<ForwardIt>::value_type> *out,
            const std::vector<sycl::event> &deps= {}) {
  auto problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};

  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;
  using MaxPair = std::pair<ForwardIt, ValueT>;

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    MaxPair p = std::make_pair(input, *input);
    reducer.combine(p);
  };

  auto reduce = [first] (MaxPair a, MaxPair b) {
    // Preserve strict total order over two equivalent
    // pointers, i.e. return the element that appears
    // in the sequence nearest to first.
    if (!(a.second < b.second) && !(b.second < a.second)) {
      if (std::distance(first, a.first) < std::distance(first, b.first))
        return a;
      else
       return b;
    }
#if __cplusplus < 202002L
    else if (a.second < b.second)
#else
    else if (std::less{}(a.second, b.second))
#endif
      return b;
    else
      return a;
  };

  MaxPair init = std::make_pair(first, *first);

  return detail::transform_reduce_impl(q, scratch_allocations, out, init,
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt, class Compare>
sycl::event
max_element(sycl::queue &q, util::allocation_group &scratch_allocations,
            ForwardIt first, ForwardIt last, Compare comp,
            std::pair<ForwardIt, typename
            std::iterator_traits<ForwardIt>::value_type> *out,
            const std::vector<sycl::event> &deps= {}) {
  auto problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};

  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;
  using MaxPair = std::pair<ForwardIt, ValueT>;

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    MaxPair p = std::make_pair(input, *input);
    reducer.combine(p);
  };

  auto reduce = [comp, first] (MaxPair a, MaxPair b) {
    // Comp used for associative containers must always
    // return false for equal values. (Effective STL, Item 21)
    // In cases where it does not (for eg. std::less_equal),
    // implementation aligns behaviour to libstdc++, returning
    // the element furthest from first.
    if (std::distance(first, a.first) < std::distance(first, b.first))
      if(comp(a.second, b.second) == false)
        return a;
      else
        return b;
    else
      if(comp(b.second, a.second) == false)
        return b;
      else
        return a;
  };

  MaxPair init = std::make_pair(first, *first);

  return detail::transform_reduce_impl(q, scratch_allocations, out, init,
                                       problem_size, kernel, reduce, deps);
}

template <class ForwardIt1, class ForwardIt2>
sycl::event equal(sycl::queue &q,
                   ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                    detail::early_exit_flag_t* out, const std::vector<sycl::event>& deps= {}) {
  std::size_t problem_size = std::distance(first1, last1);
  if(problem_size == 0)
    return sycl::event{};
  auto evt = detail ::early_exit_for_each(q, problem_size, out,
                                          [=](sycl::id<1> idx) -> bool {
                                            auto it1 = first1;
                                            auto it2 = first2;
                                            std::advance(it1, idx[0]);
                                            std::advance(it2, idx[0]);
                                            return !(*it1 == *it2);
                                          }, deps);
  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template <class ForwardIt1, class ForwardIt2, class BinaryPred>
sycl::event equal(sycl::queue &q,
                   ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                   BinaryPred p, detail::early_exit_flag_t* out,
                   const std::vector<sycl::event>& deps = {}) {
  std::size_t problem_size = std::distance(first1, last1);
  if(problem_size == 0)
    return sycl::event{};
  auto evt = detail ::early_exit_for_each(q, problem_size, out,
                                          [=](sycl::id<1> idx) -> bool {
                                            auto it1 = first1;
                                            auto it2 = first2;
                                            std::advance(it1, idx[0]);
                                            std::advance(it2, idx[0]);
                                            return !p(*it1, *it2);
                                          }, deps);
  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template <class ForwardIt1, class ForwardIt2>
sycl::event equal(sycl::queue &q,
                   ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                   ForwardIt2 last2, detail::early_exit_flag_t* out,
                   const std::vector<sycl::event>& deps= {}) {
  std::size_t size1 = std::distance(first1, last1);
  std::size_t size2 = std::distance(first2, last2);
  if (size1 != size2 || size1 == 0)
    return sycl::event{};

  auto evt = detail ::early_exit_for_each(q, size1, out,
                                          [=](sycl::id<1> idx) -> bool {
                                            auto it1 = first1;
                                            auto it2 = first2;
                                            std::advance(it1, idx[0]);
                                            std::advance(it2, idx[0]);
                                            return !(*it1 == *it2);
                                          }, deps);

  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template <class ForwardIt1, class ForwardIt2, class BinaryPred>
sycl::event equal(sycl::queue &q,
                   ForwardIt1 first1, ForwardIt1 last1, ForwardIt2 first2,
                   ForwardIt2 last2, BinaryPred p, detail::early_exit_flag_t* out,
                   const std::vector<sycl::event>& deps = {}) {
  std::size_t size1 = std::distance(first1, last1);
  std::size_t size2 = std::distance(first2, last2);
  if(size1 != size2 || size1 == 0)
    return sycl::event{};

  auto evt = detail ::early_exit_for_each(q, size1, out,
                                          [=](sycl::id<1> idx) -> bool {
                                            auto it1 = first1;
                                            auto it2 = first2;
                                            std::advance(it1, idx[0]);
                                            std::advance(it2, idx[0]);
                                            return !p(*it1, *it2);
                                          }, deps);

  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template <class RandomIt, class Compare>
sycl::event sort(sycl::queue &q, RandomIt first, RandomIt last,
                 Compare comp = std::less<>{},
                 const std::vector<sycl::event>& deps = {}) {
  std::size_t problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};

  return sorting::bitonic_sort(q, first, last, comp, deps);
}

template <class ForwardIt>
sycl::event is_sorted(sycl::queue &q, ForwardIt first, ForwardIt last,
                      detail::early_exit_flag_t* out,
                      const std::vector<sycl::event>& deps = {}) {
  std::size_t problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};

  auto evt = detail::early_exit_for_each(q, problem_size, out,
                                     [=](sycl::id<1> idx) -> bool {
                                       auto it = first;
                                       std::advance(it, idx[0]);
                                       if(it != last-1) {
                                        auto next = std::next(it, 1);
#if __cplusplus < 202002L
                                        return (*next < *it);
#else
                                        return std::less{}(*next, *it);
#endif
                                       }
                                       return false;
                                     }, deps);
  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template <class ForwardIt, class Compare>
sycl::event is_sorted(sycl::queue &q, ForwardIt first, ForwardIt last,
                      detail::early_exit_flag_t* out,
                      Compare comp,
                      const std::vector<sycl::event>& deps = {}) {
  std::size_t problem_size = std::distance(first, last);
  if(problem_size == 0)
    return sycl::event{};

  auto evt = detail::early_exit_for_each(q, problem_size, out,
                                     [=](sycl::id<1> idx) -> bool {
                                       auto it = std::next(first, idx[0]);
                                       if(it != last-1) {
                                        auto next = std::next(it, 1);
                                        return comp(*next, *it);
                                       }
                                       return false;
                                     }, deps);
  return q.single_task(evt, [=](){
    *out = static_cast<detail::early_exit_flag_t>(!(*out));
  });
}

template<class ForwardIt>
sycl::event is_sorted_until(sycl::queue &q, util::allocation_group &scratch_allocations,
                            ForwardIt first, ForwardIt last,
                            typename std::iterator_traits<ForwardIt>::difference_type *out,
                            const std::vector<sycl::event>& deps = {}) {
  if(first == last)
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [=] (ForwardIt input) {
    auto next = std::next(input, 1);
#if __cplusplus < 202002L
      if((*next < *input))
#else
      if(std::less{}(*next, *input))
#endif
        return std::distance(first, input);

    return problem_size;
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    reducer.combine(transform(input));
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out, std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}

template<class ForwardIt, class Compare>
sycl::event is_sorted_until(sycl::queue &q, util::allocation_group &scratch_allocations,
                            ForwardIt first, ForwardIt last,
                            typename std::iterator_traits<ForwardIt>::difference_type *out,
                            Compare comp,
                            const std::vector<sycl::event>& deps = {}) {
  if(first == last)
    return sycl::event{};

  using DiffT = typename std::iterator_traits<ForwardIt>::difference_type;
  DiffT problem_size = std::distance(first, last);

  auto transform = [=] (ForwardIt input) {
    auto next = std::next(input, 1);
      if(comp(*next, *input))
        return std::distance(first, input);

    return problem_size;
  };

  auto kernel = [=](sycl::id<1> idx, auto& reducer) {
    auto input = first;
    std::advance(input, idx[0]);
    reducer.combine(transform(input));
  };

  auto reduce = sycl::minimum<DiffT>{};

  return detail::transform_reduce_impl(q, scratch_allocations, out, std::numeric_limits<DiffT>::max(),
                                       problem_size, kernel, reduce, deps);
}

template< class ForwardIt1, class ForwardIt2,
          class ForwardIt3, class Compare >
sycl::event merge(sycl::queue& q,
                  util::allocation_group &scratch_allocations,
                  ForwardIt1 first1, ForwardIt1 last1,
                  ForwardIt2 first2, ForwardIt2 last2,
                  ForwardIt3 d_first, Compare comp = std::less<>{},
                  const std::vector<sycl::event>& deps = {}) {

  std::size_t size1 =  std::distance(first1, last1);
  std::size_t size2 =  std::distance(first2, last2);

  if(size1 == 0)
    return copy(q, first2, last2, d_first);
  if(size2 == 0)
    return copy(q, first1, last1, d_first);

  std::size_t problem_size = size1 + size2;
  if(problem_size == 0)
    return sycl::event{};

  if (q.get_device().get_backend() == sycl::backend::omp)
    return merging::segmented_merge(q, first1, last1, first2, last2, d_first,
                                    comp, 128, deps);
  else
    return merging::hierarchical_hybrid_merge(q, scratch_allocations, first1,
                                              last1, first2, last2, d_first,
                                              comp, 128, deps);
}

}

#endif
